/*
 * Copyright 1993-2006 NVIDIA Corporation.  All rights reserved.
 *
*/

/*
Matrix multiplication: P = M * N.
 * Device code.
*/

/* ! Simple test kernel for device functionality */
/* ! @param g_idata  input data in global memory */
/* ! @param g_odata  output data in global memory */


__kernel void matrixMul(__global float * P, __global const float * M,
                        __global const float * N, const int Mh, 
                       const int Mw, const int Nw, 
                       const int block_size) {

   const int bx = get_group_id(0);
   const int by = get_group_id(1);
   const int tx = get_local_id(0);
   const int ty = get_local_id(1);
   float Psub = 0;
   int i;
   int row, col;

   // ==================================================================
   // Solution part 4 
   // Determine the output index of each thread.
   // Compute the dot product of one row of M and one column of N
   // for each thread.
   // Write the computed value to matrix P at the correct output index
   // ==================================================================

   // Calculate the row index of the P element and M
   row = by * block_size + ty;
   
   // Calculate the column index of P and N
   col = bx * block_size + tx;

   // Destination matrix index 
   // Set indexP to reference the output element of this thread 
	
   // each thread computes one element of the work group's sub-matrix

   // For each index from [0, Width of M] 
   for (i=0; i<Mw; i ++ ) 
   {
       float MElement = M[row * Mw + i];
       float NElement = N[i * Mw + col];
       
       Psub += (MElement * NElement);
   }
   // Copy result into corresponding element of C

   P[row * Mw + col] = Psub;
   // End of solution part 4 ===========================================
}

